{
 "cells": [
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "# 代码生成：静态初始化"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 1,
   "metadata": {},
   "outputs": [],
   "source": [
    "import set_env"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": null,
   "metadata": {},
   "outputs": [],
   "source": [
    "import tvm\n",
    "from tvm import te\n",
    "import ctypes\n",
    "import numpy as np\n",
    "\n",
    "\n",
    "def test_static_callback():\n",
    "    dtype = \"int64\"\n",
    "    n = te.size_var(\"n\")\n",
    "    Ab = tvm.tir.decl_buffer((n,), dtype)\n",
    "    i = te.size_var(\"i\")\n",
    "    ib = tvm.tir.ir_builder.create()\n",
    "    A = ib.buffer_ptr(Ab)\n",
    "    cp = te.thread_axis((0, 1), \"cop\")\n",
    "    finit = tvm.tir.StringImm(\"TVMBackendRunOnce\")\n",
    "    ib.scope_attr(cp, \"coproc_uop_scope\", finit)\n",
    "    with ib.for_range(0, n, \"i\", kind=\"parallel\") as i:\n",
    "        A[i] = A[i] + 1\n",
    "    stmt = ib.get()\n",
    "\n",
    "    mod = tvm.IRModule.from_expr(tvm.tir.PrimFunc([Ab], stmt).with_attr(\"global_symbol\", \"ramp\"))\n",
    "    f = tvm.driver.build(mod, target=\"llvm\")\n",
    "    a = tvm.nd.array(np.zeros(10, dtype=dtype))\n",
    "    f(a)\n",
    "    f(a)\n",
    "    np.testing.assert_equal(a.numpy(), np.ones(a.shape[0]))\n",
    "\n",
    "\n",
    "def test_static_init():\n",
    "    dtype = \"int64\"\n",
    "    n = te.size_var(\"n\")\n",
    "    Ab = tvm.tir.decl_buffer((n,), dtype)\n",
    "    i = te.size_var(\"i\")\n",
    "    ib = tvm.tir.ir_builder.create()\n",
    "    handle = tvm.tir.call_intrin(\"handle\", \"tir.tvm_static_handle\")\n",
    "    ib.emit(tvm.tir.call_packed(\"test_static_callback\", handle, Ab))\n",
    "\n",
    "    @tvm.register_func(\"test_static_callback\")\n",
    "    def test_cb(sh, A):\n",
    "        assert isinstance(sh, ctypes.c_void_p)\n",
    "        return sh\n",
    "\n",
    "    stmt = ib.get()\n",
    "    mod = tvm.IRModule.from_expr(tvm.tir.PrimFunc([Ab], stmt).with_attr(\"global_symbol\", \"ramp\"))\n",
    "    f = tvm.driver.build(mod, target=\"llvm\")\n",
    "    a = tvm.nd.array(np.zeros(10, dtype=dtype))\n",
    "    f(a)\n",
    "\n"
   ]
  }
 ],
 "metadata": {
  "kernelspec": {
   "display_name": "py312x",
   "language": "python",
   "name": "python3"
  },
  "language_info": {
   "codemirror_mode": {
    "name": "ipython",
    "version": 3
   },
   "file_extension": ".py",
   "mimetype": "text/x-python",
   "name": "python",
   "nbconvert_exporter": "python",
   "pygments_lexer": "ipython3",
   "version": "3.12.3"
  }
 },
 "nbformat": 4,
 "nbformat_minor": 2
}
